分析Unity在移动设备的GPU内存机制（iOS篇）

开发手机游戏时，常听到身边的人传授经验：“CPU和GPU是共享一份内存的”，但这句经验到底具体指的是什么，仿佛总得不到细节精确的回答。

因此，本文尝试以一张贴图纹理的虚拟内存占用为例，就以下问题进行分析和解答：

为清晰表达避免概念混淆，本文采取以下术语： 物理内存（Physical Memory）：具体的存储硬件，各种SDRAM，比如LPDDR是移动设备常用的一种低功耗SDRAM。 虚拟内存（Virtual Memory）：对物理内存的一种逻辑映射。 系统内存（System Memory/Primary Memory）：CPU能读写的虚拟内存。 显存（Graphics Memory）：GPU能读写的虚拟内存。

另外，外存（External storage）：外部存储，“硬盘”，在移动设备一般是Flash。

如下4图[1][2]所示，iPhone6只有A8里拥有一块物理内存（1GB LPDDR3 RAM），且CPU/GPU晶片中并无物理内存（SDRAM），只有物理内存的接口（SDRAM Interface）。 且A8采取PoP封装（Package on Package），即将CPU/GPU晶片和物理内存竖直排列于A8芯片中，将CPU/GPU晶片移除后，在下一层露出了它俩共用的一块物理内存。 注，晶片中有高速Cache缓存，类型为SRAM。

其他iOS设备，iPhone、iPad等，亦如此，硬件层面，它们的物理内存都为统一内存（Unified Memory）架构，即主存和显存都位于同样的物理内存硬件中。

而桌面电脑一般是分离物理内存（Discrete Memory）架构。

自2013年的AppleA7（iPhone 5s）起iOS设备便支持Metal[3]，考虑当下（2018）的市场份额，故只讨论支持Metal的情况，而不讨论iOS上OpenGLES的情况。

系统层面，Metal支持主存显存同时访问同一块虚拟内存，即MTLBuffer的options为MTLStorageModeShared[4,5,6]，此情况已无主存显存之分，Shared模式是Buffer（比如顶点缓存、索引缓存）的默认创建模式，在iOS中Shared也是纹理缓存的默认创建模式。

此时对该虚拟内存的修改，会同时反馈到CPU和GPU上，除非CPU准备好Buffer的内容后不再修改，但一旦CPU对Buffer进行了二次修改，为避免和GPU的访问冲突，需要有一定的同步机制，比如三重缓冲（Tripple Buffering）[7]。 Pirvate模式为GPU单独访问的虚拟内存，主要用于RenderTexture等情况[9]，并非当前重点。

虽然图形API机制如此，但不同引擎内部实现大相径庭，保守起见，具体结论应以引擎具体逻辑为准。 先以纹理为例，Unity在iOS+Metal上从纹理文件存储到最终纹理显存，其二进制流的完整流程是怎样的？ 人肉阅读分析Unity源码是耗时且可能不准确的。结合Profiler等工具进行分析，会省时精确，事半功倍。这样也可顺带对Profile工具的综合应用进行介绍。所以下面，先假设我们不知道Metal的机制，试从现象推断出原因。

先创建一个名为GFXMemory的测试demo，分别有3张分辨率足够大的4096x4096的纹理贴图，格式分别设为RGBA32、RGB24、ASTC5x5，通过运行时点击对应的区域，才单独加载对应贴图，显示在屏幕中。

准备做Profile测试先查证以下问题： 由于3张纹理分辨率非常大且开启Mipmaps，其内存占用理应是期待纹理虚拟内存 = 85.33MB + 64.00MB + 13.65MB = 162.98MB，如果最终内存稳定后，本进程的虚拟内存占用约为进程内存 ~= 启动内存 + 已加载纹理内存，即可证实纹理虚拟内存占用的确只有一份，否则如果进程虚拟内存约为进程内存 ~= 启动内存 + 2*已加载纹理内存，即可证实主存、显存各持一份纹理贴图。

Unity版本为2017.4.8f1、XCode版本为10.1、运行设备为iPhone6s。 先用Unity以Development Build进行XCode工程导出，Development Build仅仅是为了能用Unity Memory Profiler进行Profile。 XCode中对Unity-iPhone工程进行Edit Scheme，并如下图开启Malloc Stack，是为了在命令行对memorygraph使用malloc_history命令查看内存创建的堆栈。

XCode中构建版本，USB连接iPhone6s并在其上运行，等待几秒钟待内存稳定后：

点击UI加载上面3张纹理后，等待几秒钟待内存稳定后：

注意上述操作都确保游戏是一次运行针对同一进程的4次抓取结果，从而确保内存地址稳定。

我们在命令行执行命令vmmap --summary ./xcode_empty.memgraph，得到加载纹理前的虚拟内存占用约为111.3MB，如下图：

上图我们应关心“DIRTY SIZE”和“SWAPPED SIZE”，前者代表已写虚存大小、后者代表已写待压缩虚存大小。iOS和一般OS不一样，不采取虚存切页（Paging）的机制，而是采取压缩内存的机制。而在iOS中所谓的内存占用（Memory Footprint）事实上是MemoryFootprint = DirtySize + CompressedSize，iOS以MemoryFootprint的大小作为Killapp的依据。注意Swapped Size是待压缩的大小，压缩后方为Compressed Size。[8]

我们再执行命令vmmap --summary ./xcode.memgraph，得到加载纹理后的虚拟内存占用约为297.8MB，如下图：

从而，加载纹理额外虚拟内存占用 = 297.9MB - 111.3MB = 186.6MB ~= 期待纹理虚拟内存占用162.98MB，而186.6MB RGBA32 >> ASTC5x5 堆内存消耗尖刺：RGB24 > RGB32 >> ASTC5x5 虚拟内存消耗则整体呈现持续增长

我们先看最左边RGBA32的CPU消耗情况，如下两图，分别为加载RGB24纹理时CPU消耗Spike的前期和后期

不需无头绪地辛苦阅读海量引擎代码，有的放矢，立刻可精确看出Unity在加载纹理时主要工作分两部分：文件加载（File::Read()）和纹理上传（UploadTexture2DData()）。 而且发现将时间线在前后期中间不管如何细分，都只出现了上面2个主要消耗，说明了只有这两个工作线程在工作，我们只需分析它们相信已足够找出纹理加载的流程。我们也发现在整个纹理加载过程中，主线程只有非常少的Update空转占用，证实纹理加载几乎是脱离主线程工作的。

文件加载函数栈看起来比较通用，先从纹理上传的函数栈看起应该会更快解决问题。可发现其关键流程如下：

通过以上比较啰嗦的分析，可以看出就算是在Metal进行纹理上传，也难免有纹理内容拷贝的过程。用[MTLDevice newTextureWithDescriptor]创建纹理对象及其指向的纹理内容空间，把FileAssetUploadInstruction的buffer数据，加以一定处理（Crunch、纹理格式转换等），最终通过[MTLTexture replaceRegion]将纹理内容数据拷贝到了驱动虚拟内存IOKit区域里。

那到底这个buffer数据到底从哪来的？当然，从上文和类名包含“File”，已经可以猜出是从外存读取得来，但不精确证实不服气，我们将注意力回到上面的文件加载调用栈。堆栈协助代码阅读，发现很简单：

那么command->buffer的内存哪里分配而来呢？

由于内存分配的CPU消耗可能很小，就算是高精度的Sampler也可能在Time Profiler里找不到，这里我们明显要求救于Allocation，如下图，我们选择“Call Trees”分类，框选在加载纹理时，内存飙升时的时段，发现132.03MB内存是在AsyncUploadManager::ManageTextureUploadRingBufferMemory()中分配给m_DataRingBuffer

文件读取的缓存应该是在堆上分配

（AsyncUploadManager::ManageTextureUploadRingBufferMemory()图略 ）

纹理上传过程中，最大的堆内存分配是分配给了AyncUploadManager.m_DataRingBuffer

通过以上种种分析，已经掌握了不少信息和关键字，找出答案已是临门一脚了：

（AsyncUploadManager::AcquireWritePtr()图略 ）

AsyncUploadManager::ScheduleAsyncRead()从m_DataRingBuffer申请纹理内容大小的内存空间，同时将指针赋值给asyncReadCommand->buffer和ftuInstr->buffer，从而文件读取线程将纹理文件内容写到asyncReadCommand->buffer指向的堆内存，渲染线程在通过ftuInstr->buffer将纹理内容从同一堆内存获取到。

至此，回答了问题2。

最后的最后，上面提到的RGB24纹理的特殊情况，为什么其虚拟内存占用大小不是64MB，而是和RGBA32一样，都是85.3MB？结合上面已知流程，分析可知，原因是Metal并不支持RGB24，在运行时都会转为RGBA32，如下：

（metal::PixelFormat图略 ）

这能从以下Time Profiler以及Allocation栈轻易证实：

Metal不支持RGB24，交给GPU使用前需要转换为RGBA32，这能从以下Time Profiler以及Allocation栈轻易证实：

（UploadTexture()中的needConversion图略 ）

Metal不支持RGB24，交给GPU使用前需要转换为RGBA32，需要消耗CPU进行一次BlitImage。

（UploadMipPyramid()图略 ）

通过Profile结果和源码，我们证实了：iOS设备中只有一块物理内存硬件，主存地址和显存地址在同一块虚存地址空间中，虚存最终的确只有一份纹理内容位于IOKit区域中，而且该纹理内容的确就是被GPU所用的纹理。 在纹理上传过程中，Unity先在堆内存申请缓存，然后将纹理文件内容读进缓存里，然后调用图形API将该该纹理内容数据拷贝到IOKit虚存中，供GPU访问。拷贝完成后缓存视乎情况从堆内存释放。 过程中，我们展示了在iOS中各种Profile工具的实际使用方法。 也介绍了一些基础的内存知识和概念。

见Github：MobileGFXMemoryTest

打算未来才做Android的Profile实验和分析报告，但通过上面的分析看来，可以大胆预测：

手机，GPU，显存，移动设备，iPhone，iPad，iOS，安卓，Android，Mobile Device，内存，共享内存，物理内存

[1]ifixit - iPhone 6 Teardown [2]Chipworks Disassembles Apple's A8 SoC [3]Metal_(API)#Supported_GPUs [4]Metal Best Practices Guide - Resource Options [5]Metal - Resource Storage Mode [6]MTLBuffer [7]Triple Buffering [8]iOS Memory Deep Dive [9]Choosing a Resource Storage Mode in iOS and tvOS [10]MTLBuffer makeTexture